package hpc.exercices.exercice3;

import static org.jocl.CL.CL_CONTEXT_PLATFORM;
import static org.jocl.CL.CL_DEVICE_TYPE_ALL;
import static org.jocl.CL.CL_MEM_COPY_HOST_PTR;
import static org.jocl.CL.CL_MEM_READ_ONLY;
import static org.jocl.CL.CL_MEM_WRITE_ONLY;
import static org.jocl.CL.CL_TRUE;
import static org.jocl.CL.clBuildProgram;
import static org.jocl.CL.clCreateBuffer;
import static org.jocl.CL.clCreateCommandQueue;
import static org.jocl.CL.clCreateContext;
import static org.jocl.CL.clCreateKernel;
import static org.jocl.CL.clCreateProgramWithSource;
import static org.jocl.CL.clEnqueueNDRangeKernel;
import static org.jocl.CL.clEnqueueReadBuffer;
import static org.jocl.CL.clGetDeviceIDs;
import static org.jocl.CL.clGetPlatformIDs;
import static org.jocl.CL.clReleaseCommandQueue;
import static org.jocl.CL.clReleaseContext;
import static org.jocl.CL.clReleaseKernel;
import static org.jocl.CL.clReleaseMemObject;
import static org.jocl.CL.clReleaseProgram;
import static org.jocl.CL.clSetKernelArg;
import hpc.exercices.Profiler;
import hpc.exercices.Utils;
import hpc.exercices.exercice2.ScanLargeArrays;

import java.io.IOException;
import java.nio.charset.Charset;
import java.util.Arrays;
import java.util.function.IntPredicate;

import org.jocl.CL;
import org.jocl.Pointer;
import org.jocl.Sizeof;
import org.jocl.cl_command_queue;
import org.jocl.cl_context;
import org.jocl.cl_context_properties;
import org.jocl.cl_device_id;
import org.jocl.cl_event;
import org.jocl.cl_kernel;
import org.jocl.cl_mem;
import org.jocl.cl_platform_id;
import org.jocl.cl_program;
import org.junit.Assert;

/**
 * Exercise 3 - Stream Compaction
 * 
 * @author Hinterleitner Michael, Heily Daniel
 *
 */

public class StreamCompaction {

	/*
	 * open cl variables
	 */
	private static cl_context context;

	private static cl_command_queue commandQueue;

	private static cl_program program;

	private static cl_device_id device;
	
	// kernels
	private static cl_kernel kernel;
	private static cl_kernel getPassedValuesParallel;
	
	// runtime
	private static long gpuTime;

	// array elements
	private static int elements = (int) Math.pow(2, 20) - 1;



	public static void main(String[] args) {

		int[] input = Utils.randomArray(elements);

		// ////////////////////////////////////////////////////////////
		// SEQUENTIAL //

		// 3 different example predicates

		IntPredicate i1 = (x) -> x < 5;
		IntPredicate i2 = (x) -> x == 5;
		IntPredicate i3 = (x) -> x < 5;

		int[] result_seq;

		long startSeqStreamCompaction = System.currentTimeMillis();
		result_seq = SequentialStreamCompaction.filter(input, i3);
		long endSeqStreamCompaction = System.currentTimeMillis();
		System.out.println("Duration for SeqStreamCompaction: "
				+ (endSeqStreamCompaction - startSeqStreamCompaction) + " ms");
		// System.out.println("Seq result = ");
		// System.out.println(Arrays.toString(result_seq));
		// ////////////////////////////////////////////////////////////

		// Parallel version
		
		// Predicate workaround for C e.g. <,> = ---> ==
		// x > 5
		char op = '<';
		int value = 5;
		
		setUp();
		createKernels();

		int[] output = new int[input.length];
		Pointer outputPointer = Pointer.to(output);
		Pointer inputPointer = Pointer.to(input);

		cl_mem inputBuffer = clCreateBuffer(context, CL_MEM_READ_ONLY
				| CL_MEM_COPY_HOST_PTR, Sizeof.cl_int * input.length,
				inputPointer, null);
		cl_mem outputBuffer = clCreateBuffer(context, CL_MEM_WRITE_ONLY
				| CL_MEM_COPY_HOST_PTR, Sizeof.cl_int * output.length,
				outputPointer, null);

		// global worksize
		long globalWorkSize[] = new long[] { input.length };

		// local worksize
		long localWorkSize[] = new long[] { 1 };

		// setting arguments for kernel
		clSetKernelArg(kernel, 0, Sizeof.cl_mem, Pointer.to(inputBuffer));
		clSetKernelArg(kernel, 1, Sizeof.cl_mem, Pointer.to(outputBuffer));
		//predicate
		clSetKernelArg(kernel, 2, Sizeof.cl_char,
				Pointer.to(new char[] { op }));
		clSetKernelArg(kernel, 3, Sizeof.cl_int,
				Pointer.to(new int[] { value }));

		cl_event kernel_event = new cl_event();

		// execute the kernel
		clEnqueueNDRangeKernel(commandQueue, kernel, 1, null, globalWorkSize,
				localWorkSize, 0, null, kernel_event);

		// wait for kernel "end" event
		CL.clWaitForEvents(1, new cl_event[] { kernel_event });

		gpuTime += Profiler.profile(kernel_event, false);

		// read the output data
		clEnqueueReadBuffer(commandQueue, outputBuffer, CL_TRUE, 0,
				output.length * Sizeof.cl_int, outputPointer, 0, null, null);

		int[] help = new int[output.length + 1];
		System.arraycopy(output, 0, help, 0, output.length);

		
		int[] scan = ScanLargeArrays.scan(help);
		
		// length +1 um auch noch Aenderung in letzten Element feststellen zu
		// koennen.. daher aber auch zuvor schon output um 1 laenger machen
		scan = Arrays.copyOfRange(scan, 0, input.length + 1);

		System.out.println("number of passed elements = " + scan[scan.length - 1]);

		Pointer scanPointer = Pointer.to(scan);

		int[] result = new int[scan[scan.length - 1]];
		Pointer resultPointer = Pointer.to(result);

		cl_mem scanBuffer = clCreateBuffer(context, CL_MEM_READ_ONLY
				| CL_MEM_COPY_HOST_PTR, Sizeof.cl_int * scan.length,
				scanPointer, null);
		cl_mem resultBuffer = clCreateBuffer(context, CL_MEM_WRITE_ONLY
				| CL_MEM_COPY_HOST_PTR, Sizeof.cl_int * result.length,
				resultPointer, null);

		// global worksize
		globalWorkSize = new long[] { input.length };

		// local worksize
		localWorkSize = new long[] { 1 };

		// setting arguments for kernel

		clSetKernelArg(getPassedValuesParallel, 0, Sizeof.cl_mem,
				Pointer.to(inputBuffer));
		clSetKernelArg(getPassedValuesParallel, 1, Sizeof.cl_mem,
				Pointer.to(resultBuffer));
		clSetKernelArg(getPassedValuesParallel, 2, Sizeof.cl_mem,
				Pointer.to(scanBuffer));
		clSetKernelArg(getPassedValuesParallel, 3, Sizeof.cl_mem,
				Pointer.to(outputBuffer));
		kernel_event = new cl_event();

		// execute the kernel
		clEnqueueNDRangeKernel(commandQueue, getPassedValuesParallel, 1, null,
				globalWorkSize, localWorkSize, 0, null, kernel_event);

		// wait for kernel "end" event
		CL.clWaitForEvents(1, new cl_event[] { kernel_event });

		gpuTime += Profiler.profile(kernel_event, false);

		// read the output data
		clEnqueueReadBuffer(commandQueue, resultBuffer, CL_TRUE, 0,
				result.length * Sizeof.cl_int, resultPointer, 0, null, null);

		// release kernel, program, and memory objects
		clReleaseMemObject(resultBuffer);
		clReleaseMemObject(inputBuffer);
		clReleaseMemObject(outputBuffer);

		shutdown();

		Assert.assertArrayEquals(result_seq, result);
		System.out.println("parallel and seq stream compaction had same result");
		System.out.println("Duration for compaction phase: " + String.format("%8.3f", gpuTime / 1e6)
				+ " ms");
	}

	private static void setUp() {
		final int platformIndex = 0;
		final long deviceType = CL_DEVICE_TYPE_ALL;
		final int deviceIndex = 0;

		// enable exceptions
		CL.setExceptionsEnabled(true);

		// obtain the number of platforms
		int numPlatformsArray[] = new int[1];
		clGetPlatformIDs(0, null, numPlatformsArray);
		int numPlatforms = numPlatformsArray[0];

		// obtain a platform ID
		cl_platform_id platforms[] = new cl_platform_id[numPlatforms];
		clGetPlatformIDs(platforms.length, platforms, null);
		cl_platform_id platform = platforms[platformIndex];

		// initialize the context properties
		cl_context_properties contextProperties = new cl_context_properties();
		contextProperties.addProperty(CL_CONTEXT_PLATFORM, platform);

		// obtain the number of devices for the platform
		int numDevicesArray[] = new int[1];
		clGetDeviceIDs(platform, deviceType, 0, null, numDevicesArray);
		int numDevices = numDevicesArray[0];

		// obtain a device ID
		cl_device_id devices[] = new cl_device_id[numDevices];
		clGetDeviceIDs(platform, deviceType, numDevices, devices, null);
		device = devices[deviceIndex];

		// create a context for the selected device
		context = clCreateContext(contextProperties, 1,
				new cl_device_id[] { device }, null, null, null);

		long properties = 0;
		properties |= CL.CL_QUEUE_PROFILING_ENABLE;
		properties |= CL.CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE;

		// create a command-queue for the selected device cl_command_queue
		commandQueue = clCreateCommandQueue(context, device, properties, null);
	}

	private static void createKernels() {

		String programSource = null;
		try {
			programSource = Utils.readFile(
					"kernelCode/exercise3_streamCompaction.cl",
					Charset.defaultCharset());

		} catch (IOException e) {
			System.err.println("kernel source code couldn't be read! Message:"
					+ e.getMessage());
		}

		// create the program from the source code
		program = clCreateProgramWithSource(context, 1,
				new String[] { programSource }, null, null);

		// build the program
		clBuildProgram(program, 0, null, null, null, null);

		// create the kernel
		kernel = clCreateKernel(program, "prepare", null);
		getPassedValuesParallel = clCreateKernel(program,
				"getPassedValuesParallel", null);
	}

	private static void shutdown() {
		clReleaseKernel(kernel);
		clReleaseKernel (getPassedValuesParallel);
		clReleaseProgram(program);
		clReleaseCommandQueue(commandQueue);
		clReleaseContext(context);
	}

}
